
#include <CL/cl.h>
#include <stdio.h>
#include <sys/time.h>
#include <CL/cl.h>
#include <unistd.h>


void my_cl_check_error(cl_int error);

#define TEST_SIZE (1024*100)


int main(int argc, char** argv) {

	///////////////////////////////////////////////////////////////////
	/////////    DEVICE INTITALIZATION    /////////////////////////////
	///////////////////////////////////////////////////////////////////


	

    cl_platform_id cl_platform;       //OpenCL platform
    cl_device_id cl_device;           //OpenCL device
    cl_context cl_gpu_context;    //OpenCL context
    cl_int cl_errnum;
    
    //searching for available opencl platforms in the system
    cl_platform_id search_platforms[64];
    cl_uint nplatforms = 0;
    printf("Searching for available OpenCL platforms...\n");
    clGetPlatformIDs( 64, search_platforms, &nplatforms );
	printf("Found %d platforms!!\n", nplatforms);
	for (int i = 0; i < nplatforms; i++) {
			char plat_name[128];
			cl_errnum = clGetPlatformInfo(search_platforms[i], CL_PLATFORM_NAME, 128, plat_name, NULL);
			my_cl_check_error(cl_errnum);
			printf("	platform %d is: %s...\n",i, plat_name);	
	}
	
		if (argc != 5) {
		printf("USAGE: -p 'plaftorm number' -d 'device number'\n");
		exit(0);
	}

	int selected_platform = atoi(argv[2]);
	printf("Choosing platform %d \n", selected_platform);	
	cl_platform = search_platforms[selected_platform];
	my_cl_check_error(cl_errnum);
	cl_device_id search_devices[64];
	cl_uint ndevices;
    clGetDeviceIDs(cl_platform, CL_DEVICE_TYPE_ALL, 64, search_devices, &ndevices);
    printf("This platform has %d devices\n", ndevices);
    for (int i = 0; i < ndevices; i++) {
    	char device_name[128];
		cl_errnum = clGetDeviceInfo(search_devices[i], CL_DEVICE_NAME, 128, device_name, NULL);
		my_cl_check_error(cl_errnum);
    	printf("	device %d is: %s\n", i, device_name); 
    }
	int selected_device = atoi(argv[4]);
	printf("choosing device: %d \n", selected_device);
	cl_device = search_devices[selected_device];
	my_cl_check_error(cl_errnum);


	cl_gpu_context = clCreateContext(0, 1, &cl_device, NULL, NULL, &cl_errnum);
	my_cl_check_error(cl_errnum);
	cl_command_queue command_queue;
	command_queue = clCreateCommandQueue(cl_gpu_context, cl_device, CL_QUEUE_PROFILING_ENABLE, &cl_errnum );
	my_cl_check_error(cl_errnum);
	
	///////////////////////////////////////////////////////////////////
	/////////   END OF DEVICE INTITALIZATION    ///////////////////////
	///////////////////////////////////////////////////////////////////
	
	
	///////////////////////////////////////////////////////////////////
	/////////////   LOAD KERNELS //////////////////////////////////////
	///////////////////////////////////////////////////////////////////
	
	
	FILE* fil = fopen("testkernel.cl","r");
	char* prog_string = (char*) malloc( 22048 );
	fread(prog_string, 1, 22048, fil);
	
	cl_program program = clCreateProgramWithSource(cl_gpu_context, 1, (const char **)&prog_string, NULL, &cl_errnum);
	my_cl_check_error(cl_errnum);
	cl_errnum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	if (cl_errnum) {
		printf("ERROR building\n");
		size_t len;
		char buffer[8048];
		// get the build log
		clGetProgramBuildInfo( program, cl_device, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
		printf("--- Build Log -- \n %s\n",buffer);
		//exit(1);
		my_cl_check_error(cl_errnum);
	}
	my_cl_check_error(cl_errnum);

	
	//test kernel
	{

		//create the input data on the host
		int data_size = TEST_SIZE*sizeof(int);
		int* host_source_frame = (int*) malloc( data_size );
		//initialize
		for (int i = 0; i < TEST_SIZE; i++) {
			host_source_frame[i] = i;
		}
		//create the input data on the device and initialize it with the host data
		cl_mem source_frame = clCreateBuffer( cl_gpu_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, data_size, host_source_frame, &cl_errnum );
		my_cl_check_error(cl_errnum);

		//create the output data on the host
		int* host_dest_frame = (int*) malloc( data_size );
		//create the output data on the device
		cl_mem dest_frame = clCreateBuffer( cl_gpu_context, CL_MEM_WRITE_ONLY, data_size, NULL, &cl_errnum );
		my_cl_check_error(cl_errnum);

		//set the parameters for the kernel call
		cl_kernel test_kernel = clCreateKernel(program, "granularity_test_kernel", &cl_errnum);
		my_cl_check_error(cl_errnum);

		clSetKernelArg( test_kernel, 0, sizeof(cl_mem), &source_frame);
		clSetKernelArg( test_kernel, 1, sizeof(cl_mem), &dest_frame);
		
        cl_event timing_event;
        printf("Checking parallel granularity of the device, look for an improvement in execution times when reaching certain granulaity\n");
        printf("| LOCAL_WORK_SIZE |  GRANULARITY | EXECUTION_TIME | FAST_BRANCH_TIME | SLOW_BRANCH_TIME | IDEAL_TIME_WITH_BRANCHING |\n");
        for (int lsize = 32; lsize < 513; lsize*=2) {
        	long worst_branch_time = 0;
			long best_branch_time = 0;
			for ( int gi = 0; gi < 13; gi ++ ) {

				int granularity_table[13] = { -1, 0, 1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024 };
				int granularity = granularity_table[gi];
				if (granularity > lsize) continue;
		        for (int i = 0; i < TEST_SIZE; i++) {
		        	if (granularity == -1) {
		        		host_source_frame[i] = (i+1)*2; //for the slow branch, all are odd
		        	}
		        	else if ( granularity== 0) {
		        		host_source_frame[i] = (i+1)*2 +1; //for the fast branch, all are even
		        	}
		        	else {
						host_source_frame[i] = i/granularity; //a way create cluster of size granularity of the same branch
					}
				}

				cl_errnum = clEnqueueWriteBuffer( command_queue, source_frame, CL_TRUE, 0, data_size, host_source_frame, 0, NULL, NULL);
			    my_cl_check_error(cl_errnum);

        		unsigned long total_time = 0;
        		//launch the kernel
				for (int ii = 0; ii < 50; ii++) {

					size_t global_work_size = { TEST_SIZE };
					size_t local_work_size = { lsize };
					cl_errnum = clEnqueueNDRangeKernel( command_queue, test_kernel, 1, NULL,  &global_work_size, &local_work_size, 0, NULL, &timing_event);
					my_cl_check_error(cl_errnum);
					cl_errnum = clFinish(command_queue);
					my_cl_check_error(cl_errnum);

					cl_errnum = clWaitForEvents(1, &timing_event);
					my_cl_check_error(cl_errnum);
					cl_ulong k_start_time, k_end_time;
					clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_START,sizeof(cl_ulong), &k_start_time, NULL);
					clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_END,sizeof(cl_ulong), &k_end_time, NULL);
					//printf("intra time of the MONO kernel: %ld nsecs\n", k_end_time - k_start_time);
					if (ii > 4) { //let the device warmup a little, do no count first kernel calls
						total_time += k_end_time - k_start_time;
					}
				}
				total_time /= 50-4;

				if (granularity == -1) {
					worst_branch_time = total_time;
				}
				else if (granularity == 0) {
					best_branch_time = total_time;
				}
				else {
					printf("| %15d | %12d | %14ld | %16ld | %16ld | %25ld |\n", lsize, granularity, total_time, best_branch_time, worst_branch_time, (best_branch_time+ worst_branch_time)/2 );
				}
			}
		}
		printf("| ALL TIMES IN NANOSECONDS |\n\n");
        /*
        cl_errnum = clEnqueueReadBuffer( command_queue, dest_frame, CL_TRUE, 0, data_size, host_dest_frame, 0, NULL, NULL);
		my_cl_check_error(cl_errnum);
		for (int i = 0; i < 16; i++) {
			printf("res %d\n", host_dest_frame[i]);
		}
			*/
			
				{
	
		
		int frame_width = 2048;
		int frame_height = 2048;
		int data_size = frame_width * frame_height;
			

		cl_mem current_frame = clCreateBuffer( cl_gpu_context, 0, data_size, NULL, &cl_errnum );
		my_cl_check_error(cl_errnum);

		cl_mem reference_frame = clCreateBuffer( cl_gpu_context, 0, data_size, NULL, &cl_errnum );
		my_cl_check_error(cl_errnum);
		
		cl_mem output_vectors = clCreateBuffer( cl_gpu_context, 0, data_size, NULL, &cl_errnum );
		my_cl_check_error(cl_errnum);

		//set the parameters for the kernel call
		cl_kernel fs_kernel = clCreateKernel(program, "test_fullsearch_lowerbound", &cl_errnum);
		my_cl_check_error(cl_errnum);
		
		unsigned int stride = frame_width;
		cl_errnum = clSetKernelArg( fs_kernel, 0, sizeof(cl_mem), &current_frame); my_cl_check_error(cl_errnum);
		cl_errnum = clSetKernelArg( fs_kernel, 1, sizeof(unsigned int), &stride); my_cl_check_error(cl_errnum);
		cl_errnum = clSetKernelArg( fs_kernel, 2, sizeof(cl_mem), &reference_frame); my_cl_check_error(cl_errnum);
		cl_errnum = clSetKernelArg( fs_kernel, 3, sizeof(unsigned int), &stride); my_cl_check_error(cl_errnum);
		cl_errnum = clSetKernelArg( fs_kernel, 4, sizeof(cl_mem), &output_vectors); my_cl_check_error(cl_errnum);
		cl_errnum = clSetKernelArg( fs_kernel, 5, sizeof(unsigned int), &stride); my_cl_check_error(cl_errnum);
	
    	printf("Executing a test to find a lower bound for full search...\n");
#define FS_BLOCK_X 64
#define FS_PIX_PER_THREAD 4
#define FS_MBLOCK_PER_CUDABLOCK ((FS_BLOCK_X * FS_PIX_PER_THREAD)/ 16)
#define FS_CHECK_POINTS 1024
		unsigned long fullsearch_time = 0;
		for (int tt = 0; tt < 100; tt++) {
			size_t global_work_size[2] = { frame_width/(16), frame_height/16 };
			size_t local_work_size[2] = { 64, 1 };
			cl_errnum = clEnqueueNDRangeKernel( command_queue, fs_kernel, 2, NULL,  global_work_size, local_work_size, 0, NULL, &timing_event);
			my_cl_check_error(cl_errnum);
			cl_errnum = clFinish(command_queue);
			my_cl_check_error(cl_errnum);


			cl_errnum = clWaitForEvents(1, &timing_event);
			my_cl_check_error(cl_errnum);
			cl_ulong k_start_time, k_end_time;
			clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_START,sizeof(cl_ulong), &k_start_time, NULL);
			clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_END,sizeof(cl_ulong), &k_end_time, NULL);
			//printf("time for fullsearch %ld \n", k_end_time - k_start_time);
			fullsearch_time += k_end_time - k_start_time;
		}
		printf("Lower bound for full search on a %d x %d frame on %d check points: %ld nsecs\n", frame_width, frame_height, FS_CHECK_POINTS, fullsearch_time/100 );
	}
			
			
	}
	


}
























void my_cl_check_error(cl_int error)
{
    static const char* errorString[] = {
        "CL_SUCCESS",
        "CL_DEVICE_NOT_FOUND",
        "CL_DEVICE_NOT_AVAILABLE",
        "CL_COMPILER_NOT_AVAILABLE",
        "CL_MEM_OBJECT_ALLOCATION_FAILURE",
        "CL_OUT_OF_RESOURCES",
        "CL_OUT_OF_HOST_MEMORY",
        "CL_PROFILING_INFO_NOT_AVAILABLE",
        "CL_MEM_COPY_OVERLAP",
        "CL_IMAGE_FORMAT_MISMATCH",
        "CL_IMAGE_FORMAT_NOT_SUPPORTED",
        "CL_BUILD_PROGRAM_FAILURE",
        "CL_MAP_FAILURE",
        "",
        "",
        "",
        "",
        "",
        "",
        "",
        "",
        "",
        "",
        "",
        "",
        "",
        "",
        "",
        "",
        "",
        "CL_INVALID_VALUE",
        "CL_INVALID_DEVICE_TYPE",
        "CL_INVALID_PLATFORM",
        "CL_INVALID_DEVICE",
        "CL_INVALID_CONTEXT",
        "CL_INVALID_QUEUE_PROPERTIES",
        "CL_INVALID_COMMAND_QUEUE",
        "CL_INVALID_HOST_PTR",
        "CL_INVALID_MEM_OBJECT",
        "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR",
        "CL_INVALID_IMAGE_SIZE",
        "CL_INVALID_SAMPLER",
        "CL_INVALID_BINARY",
        "CL_INVALID_BUILD_OPTIONS",
        "CL_INVALID_PROGRAM",
        "CL_INVALID_PROGRAM_EXECUTABLE",
        "CL_INVALID_KERNEL_NAME",
        "CL_INVALID_KERNEL_DEFINITION",
        "CL_INVALID_KERNEL",
        "CL_INVALID_ARG_INDEX",
        "CL_INVALID_ARG_VALUE",
        "CL_INVALID_ARG_SIZE",
        "CL_INVALID_KERNEL_ARGS",
        "CL_INVALID_WORK_DIMENSION",
        "CL_INVALID_WORK_GROUP_SIZE",
        "CL_INVALID_WORK_ITEM_SIZE",
        "CL_INVALID_GLOBAL_OFFSET",
        "CL_INVALID_EVENT_WAIT_LIST",
        "CL_INVALID_EVENT",
        "CL_INVALID_OPERATION",
        "CL_INVALID_GL_OBJECT",
        "CL_INVALID_BUFFER_SIZE",
        "CL_INVALID_MIP_LEVEL",
        "CL_INVALID_GLOBAL_WORK_SIZE",
    };

    const int errorCount = sizeof(errorString) / sizeof(errorString[0]);

    const int index = -error;

    if (index > 0 && index < errorCount) {
    	printf("Error in openCL: %s \n", errorString[index]);
    	exit(0);
    }

}
